#include "book.h"
#include "cuPrintf.cu"
#define numThreadsPerBlock 2 
#define topN  10
#define multfactor 1


struct cuRadix {
int score;
int docNumber;

};

__device__ int prefix_scan (cuRadix* d_array )
{
	uint tid = threadIdx.x;
	uint size = blockDim.x;
	for( int i = 1; i < size; i *= 2)
	{
		cuRadix dtemp_cr;
		if ( tid >= i )
		{
			dtemp_cr.score = d_array[tid-i].score;			
		}
		__syncthreads();

		if ( tid >= i )
		{
			d_array[tid].score += dtemp_cr.score ;
		}	

		__syncthreads();

	}
	return d_array[tid].score;
}

__device__ void radix_sort (cuRadix* d_array ,int d_insize ) 
{

	for (int bit_pos = 0 ; bit_pos < 32 ; bit_pos ++ )
	{
		uint tid = threadIdx.x;
		uint size = blockDim.x;
		if (tid < d_insize)
		{
			uint value = d_array[tid].score;
			int DN	= d_array[tid].docNumber;
			uint bitvalue = (value >> bit_pos) & 1;
			d_array[tid].score = bitvalue;
			
			__syncthreads();

			uint scan_result = prefix_scan(d_array);
			uint oneBitSum  = d_array[size-1].score;
			uint zeroBitSum  = size - oneBitSum;
			__syncthreads();

			if ( bitvalue )
			{
	      			d_array[scan_result-1 + zeroBitSum].score = value;
				d_array[scan_result-1 + zeroBitSum].docNumber = DN;
			}
	    		else
			{
	      			d_array[tid - scan_result].score = value;
				d_array[tid - scan_result].docNumber = DN;
			}
	}	
		else
			cuPrintf("My tid is more than the array size \n");
	
	}
}

//Obtained from Nvdia document
__device__ void prescan(cuRadix *g_odata, cuRadix* g_idata, int n)
{
     __shared__ int temp[1024];// allocated on invocation
    int thid = threadIdx.x;
    int offset = 1;
    temp[2*thid] = g_idata[2*thid].score; // load input into shared memory
    temp[2*thid+1] = g_idata[2*thid+1].score;
     for ( int d = n>>1; d > 0; d >>= 1 ) // build sum in place up the tree
    {
        __syncthreads();
        if ( thid < d )
        {
            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            temp[bi] += temp[ai];
        }
        offset *= 2;
    }
    if ( thid == 0 )
    {
        temp[n - 1] = 0;
    } // clear the last element
   for ( int d = 1; d < n; d *= 2 ) // traverse down tree & build scan
    {
        offset >>= 1;
        __syncthreads();
        if ( thid < d )
        {
            int ai = offset*(2*thid+1)-1;
            int bi = offset*(2*thid+2)-1;
            int t = temp[ai];
            temp[ai] = temp[bi];
            temp[bi] += t;
        }
    }
    __syncthreads();
    g_odata[2*thid].score = temp[2*thid]; // write results to device memory
    g_odata[2*thid+1].score = temp[2*thid+1];
    cuPrintf("d_array[%d] : %d \n",2*thid , temp[2*thid]);
    cuPrintf("d_array[%d] : %d \n",2*thid+1 , temp[2*thid]+1);
}

__device__ void radix_sort_1024 (cuRadix* d_array ,int d_insize ) 
{

	for (int bit_pos = 0 ; bit_pos < 1 ; bit_pos ++ )
	{
		cuPrintf("NEXT ROUND\n");		
		uint tid = threadIdx.x;
		uint size = blockDim.x;
		if (tid < d_insize)
		{
			uint value[4];
			int DN[4];
			uint bitvalue [4];

			for (int i =0 ; i <4 ; i ++)
			{
				value[i] = d_array[tid + ( i * size) ].score;
				DN[i]= d_array[tid + ( i * size)].docNumber;
				bitvalue[i] = (value[i] >> bit_pos) & 1;
				d_array[tid + ( i * size)].score = bitvalue[i];
				//cuPrintf("the index  : %d actual value : %d and the bitvalue :  % d \n", tid + ( i * size) ,value[i], bitvalue[i]);
			}
		
			
			__syncthreads();
			
			int scan_result [4];
			uint oneBitSum ;
			uint zeroBitSum; 
			
			prescan(d_array, d_array, d_insize);
			__syncthreads();
			oneBitSum  = d_array[d_insize-1].score;
			zeroBitSum  = d_insize - oneBitSum;
			cuPrintf("The one bit sum is %d and zerobit sum is %d \n",oneBitSum ,zeroBitSum);
			for (int i =0 ; i <4 ; i ++)
			{
 				scan_result[i] = d_array[tid + ( i * size) ].score;

				//cuPrintf("scanresult[%d] : %d , d_index : %d \n" ,i, scan_result[i],tid + ( i * size) );
			}
			__syncthreads();
			for (int i =0 ; i <4 ; i ++)	
			{
				if ( bitvalue )
				{      			
						d_array[scan_result[i]-1 + zeroBitSum].score = value[i];
						d_array[scan_result[i]-1 + zeroBitSum].docNumber = DN[i];
				}
		    		else
				{
							
						d_array[(tid + ( i * size) )- scan_result[i]].score = value[i];
						d_array[( tid + ( i * size) ) - scan_result[i]].docNumber = DN[i];
				}
				
			__syncthreads();
			}
		}	
		else
			cuPrintf("My tid is more than the array size \n");
	
	}


}








__global__ void topNe (cuRadix * d_input ,int size )
{
	extern __shared__ cuRadix d_array[]  ;
	int tid = threadIdx.x;

	if(tid==0)
	{
		cuPrintf("I am going to radix_sort\n");
	}
	radix_sort (d_input,size );
	__syncthreads();
}



__global__ void topNe_1024 (cuRadix * d_input ,int size )
{
	extern __shared__ cuRadix d_array[]  ;
	int tid = threadIdx.x;

	if(tid==0)
	{
		cuPrintf("I am going to radix_sort_1024\n");
	}
	radix_sort_1024 (d_input,size );
	__syncthreads();
}


int main (void) {

	cuRadix *h_input;
	cuRadix *d_input;
	int mem_size;
	mem_size = (int)numThreadsPerBlock *4;
	cuRadix *h_output ;  


 	
	//Allocate Input array on the host
  	h_input = (cuRadix *) malloc (mem_size * sizeof(cuRadix));
  
	//allocate the input array
  	HANDLE_ERROR(cudaMalloc((void **)&d_input,mem_size *sizeof(cuRadix)  ));
 	
 	//allocate the output array on host
	h_output = (cuRadix *) malloc (mem_size * sizeof(cuRadix));

	//Initialize the host input array
	int i ;
	for(i=0;i < mem_size; i ++ )
	{
		h_input[i].score = (i + 11)%4;
		h_input[i].docNumber = i;
	} 
  
  	//Copy the input arry from host ot device
	HANDLE_ERROR(cudaMemcpy(d_input,h_input, sizeof(cuRadix) * mem_size,cudaMemcpyHostToDevice));

	cudaPrintfInit();
	int threads_needed;
	//if (mem_size <= numThreadsPerBlock)
	if (mem_size == 8)
	{
		threads_needed = numThreadsPerBlock;
		topNe <<< 1, threads_needed , mem_size  >>> (d_input,mem_size);
		
	}
	else if (mem_size < numThreadsPerBlock)
	{
		threads_needed = mem_size; 
		topNe <<< 1, threads_needed , mem_size  >>> (d_input,mem_size);
	}

	
	cudaPrintfDisplay(stdout, true);
	cudaPrintfEnd();

	HANDLE_ERROR (cudaMemcpy (h_output, d_input, sizeof(cuRadix) * mem_size,cudaMemcpyDeviceToHost )  ) ;	
	
	for (i = 0 ; i < mem_size ; i ++ )
	{
		printf ("the score : %d and the doc number :  %d \n",h_output[i].score , h_output[i].docNumber );
	}
	cudaFree(d_input);
	free(h_input);
	free(h_output);
}
